原子操作的用法广泛和多样，以至于不可能在本书中提供每种用法的示例。我们展示了两个具有代表性的例子，它们具有广泛的适用性:\par
	
\begin{enumerate}
	\item 计算直方图
	\item 实现设备范围内的同步
\end{enumerate}

\hspace*{\fill} \par %插入空行
\textbf{计算直方图}

图19-18中的代码演示了如何使用自由原子操作和工作组栅栏来计算直方图。内核被栅栏分为三个阶段，每个阶段都有自己的原子性需求。这个栅栏既是一个同步点，又是一个获取-释放栅栏——这确保了一个阶段中的任何读和写，对后面阶段中工作组的所有工作项都是可见的。\par

第一阶段将某个工作组本地内存的内容设置为零。每个工作组中的工作项不能按照设计争用条件更新工作组本地内存中的位置，并且不需要原子性。\par

第二阶段将部分直方图结果累加到本地内存中。同一个工作组中的工作项可以更新工作组本地内存中的相同位置，但是同步可以推迟到该阶段的结束——可以使用memory\_order::relaxed和memory\_scope::work\_group来满足原子性需求。\par

第三阶段将部分直方图结果存储在全局内存中。保证同工作组中的工作项从工作组本地内存中的位置读取，但是会更新全局内存中的相同位置——不再需要工作组本地内存的原子性，可以像以前一样使用memory\_order::relaxed和memory\_scope::system来满足全局内存的原子性要求。\par

\hspace*{\fill} \par %插入空行
图19-18 使用不同内存空间中的原子操作计算直方图
\begin{lstlisting}[caption={}]
// Define shorthand aliases for the types of atomic needed by this kernel 
template <typename T>
using local_atomic_ref = atomic_ref<
	T,
	memory_order::relaxed, 
	memory_scope::work_group,
	access::address_space::local_space>; 
	
template <typename T>
using global_atomic_ref = atomic_ref<
	T,
	memory_order::relaxed, 
	memory_scope::system,
	access::address_space::global_space>; 
	
Q.submit([&](handler& h) {
	auto local = 1ocal_accessor<uint32_t, 1>{B, h};
	h.parallel_for(
		nd_range<1>{num_groups * num_items, num_items},[=](nd_item<1> it){ 
			// Phase 1:Work-items co-operate to zero local memory
			for (int32_t b=it.get_local_id(0); b <B; b+=it.get_local_range(0)){
				local[b]=0;
			}
		
			it.barrier();// Wait for all to be zeroed
		
			// Phase 2:Work-groups each compute a chunk of the input 
			//Work-items co-operate to compute histogram in local memory 
			auto grp=it.get_group();
		
			const auto [group_start, group_end] = distribute_range (grp, N); 
			for (int i = group_start + it.get_local_id(0); i < group_end; 
				i +=it.get_local_range(0)){
			
				int32_t b = input[i] % B;
				local_atomic_ref<uint32_t>(local[b])++;
			}
			it.barrier(); // Wait for all local histogram updates to complete 
			
			// Phase 3: Work-items co-operate to update global memory
			for (int32_t b = it.get_local_id(0); b < B; b +=it.get_local_range(0)){
				
				global_atomic_ref<uint32_t>(histogram[b]) += local[b];
			}
		});
	}).wait();
\end{lstlisting}

\hspace*{\fill} \par %插入空行
\textbf{实现设备范围内的同步}

第4章警告不要编写跨工作组同步工作项的内核。然而，我们完全期望本章的几个示例能在原子操作之上实现对设备范围的同步。\par

\begin{tcolorbox}[colback=red!5!white,colframe=red!75!black]
设备范围内的同步目前是不可移植的，最好留给专业开发者使用。语言的未来版本将解决这个问题。
\end{tcolorbox}

本节中讨论的代码是危险的，不应该期望在所有设备上工作，因为在调度和并发保证方面可能存在差异。原子操作提供的内存序保证与前进进度保证是正交的；在编写本文时，SYCL和DPC++中的工作组调度完全是由实现定义的。讨论执行模型和调度保证所需的概念和术语是目前学术研究领域的热门，DPC++的未来版本预计将在此工作的基础上提供额外的调度查询和控制。\par

图19-19展示了一个设备级锁存器(一次性栅栏)的简单实现，图19-20展示了一个简单的使用示例。每个工作组选择一个单独的工作项，以发出该组到达的信号，并使用自旋环等待其他组，而其他工作项使用工作组栅栏等待所选的工作项。正是这种自旋环使得设备范围的同步不安全，当工作组还没有开始执行，或者当前执行的工作组没有平衡地调度，代码可能会死锁。\par

\begin{tcolorbox}[colback=red!5!white,colframe=red!75!black]
如果没有进度保证，单独依赖内存序来实现同步原语可能会导致死锁!
\end{tcolorbox}

要使代码正确工作，必须满足以下三个条件:\par

\begin{enumerate}
	\item 为了保证生成正确的内存栅栏，原子操作必须使用严格的内存序。
	\item ND-Range中的每个工作组必须能够执行，避免单个工作组在循环中自旋，从而使尚未增加计数器的工作组挨饿。
	\item 设备必须能够同时执行ND-Range中的所有工作组，以确保ND-Range中的所有工作组最终达到门闩处。
\end{enumerate}

\hspace*{\fill} \par %插入空行
图19-19 原子引用的顶部构建一个简单的设备级闩锁
\begin{lstlisting}[caption={}]
struct device_latch {
	using memory_order = intel::memory_order;
	using memory_scope = intel::memory_scope;
	
	explicit device_latch(size_t num_groups) :
			counter(0), expected(num_groups) {}
			
	template <int Dimensions>
	void arrive_and_wait(nd_item<Dimensions>& it) {
		it.barrier();
		// Elect one work-item per work-group to be involved
		// in the synchronization
		// All other work-items wait at the barrier after the branch
		if (it.get_local_linear_id() == 0) {
			atomic_ref<
				size_t,
				memory_order::acq_rel,
				memory_scope::device,
				access::address_space::global_space> atomic_counter(counter);
				
			// Signal arrival at the barrier
			// Previous writes should be visible to 
			// all work-items on the device
			atomic_counter++;
			
			// Wait for all work-groups to arrive
			// Synchronize with previous releases by
			// all work-items on the device
			while (atomic_counter.load() != expected) {}
		}
		it.barrier();
	}
	size_t counter;
	size_t expected;
};
\end{lstlisting}

\hspace*{\fill} \par %插入空行
图19-20 使用图19-19中的设备级闩锁
\begin{lstlisting}[caption={}]
// Allocate a one-time-use device_latch in USM
void* ptr = sycl::malloc_shared(sizeof(device_latch), Q);
device_latch* latch = new (ptr) device_latch(num_groups);
Q.submit([&](handler& h) {
	h.parallel_for(R, [=](nd_item<1> it) {
		// Every work-item writes a 1 to its location
		data[it.get_global_linear_id()] = 1;
		
		// Every work-item waits for all writes
		latch->arrive_and_wait(it);
		
		// Every work-item sums the values it can see
		size_t sum = 0;
		for (int i = 0; i < num_groups * items_per_group; ++i) {
			sum += data[i];
		}
		sums[it.get_global_linear_id()] = sum;
	});
}).wait();
free(ptr, Q);
\end{lstlisting}

虽然这段代码不能保证可移植性，但强调两个关键点:1)DPC++具有足够的表达能力，可以实现特定于设备的调优，有时需要以牺牲可移植性为代价；2)DPC++包含了实现高级同步例程所需的构建块，这可能会包含在语言的未来版本中。\par









































